#include <stdio.h>
#include <cuda_runtime.h>


__global__ void kernel(float *da, int n) {

    int sleepMs = 30;

    int a = 0;
    for (int i = 0; i < sleepMs * 1000 * 1000; ++i) {
        a++;
    }

    printf("\nkernel1 执行\n\n");
}

__global__ void kerne2(float *da, int n) {

    int sleepMs = 20;
    int a = 0;
    for (int i = 0; i < sleepMs * 1000 * 1000; ++i) {
        a++;
    }

    printf("\nkernel2 执行\n\n");
}

__global__ void kerne3(float *da, int n) {

    int sleepMs = 30;
    int a = 0;
    for (int i = 0; i < sleepMs * 1000 * 1000; ++i) {
        a++;
    }

    printf("\nkernel3 执行\n\n");
}


int main() {

    // host数据 
    float ha[16]{0};
    ha[0] = 10.5;
    ha[3] = 5.1;

    // device数据 
    float *da;
    cudaMalloc((void **)&da, sizeof(float) * 16);

    // h2d 
    cudaMemcpy(da, ha, sizeof(float) * 16, cudaMemcpyHostToDevice);  // 同步操作 

    // 创建stream
    cudaStream_t stream1;
    cudaStream_t stream3;
    
    // cudaStreamCreate(&stream1);
    // cudaStreamCreate(&stream2);

    cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);  // 指定非空流对空流（空流就是默认流）的默认阻塞失效
    cudaStreamCreateWithFlags(&stream3, cudaStreamNonBlocking);

    // 启动kernel，正常是执行顺序是kernel1、kernel3
    // kernel<<<1, 1, 0, stream1>>>(da, 16);
    // // kerne2<<<1, 1>>>(da, 16);  // 默认流会阻塞其他流，所以不用默认流 
    // kerne3<<<1, 1, 0, stream2>>>(da, 16);


    // 通过流同步来调整kernel执行顺序，执行完kernel3，再执行kernel1
    cudaEvent_t k1_k3;  // stream3 等 stream1  
    cudaEventCreate(&k1_k3);

    // 正确顺序：先 Record，后 Wait
    // 错误顺序：先 Wait，后 Record
    kernel<<<1, 1, 0, stream1>>>(da, 16);
    cudaEventRecord(k1_k3, stream1);  // 向stream1中入队一个事件类型的task 

    cudaStreamWaitEvent(stream3, k1_k3, 0);  // 向stream3中入队一个等待类型的task，参数0为保留值 
    kerne3<<<1, 1, 0, stream3>>>(da, 16);
    

    

    

    // 等待所有流执行完  
    // cudaDeviceSynchronize();
    cudaStreamSynchronize(stream1);
    cudaStreamSynchronize(stream3);


    return 0;
}